/*!
  \file
  File to implement the image inverter for HHVolumeGPU
*/

#include "invertimagegpu.hpp"

// ==================================

void InvertImageGPU::Invert( SciGPU::HHV::HHVolume<long,2>& dst,
                             const SciGPU::HHV::HHVolume<long,2>& src ) const {
  SciGPU::HHV::HHVolumeGPU<long,2>::convertor srcConvert( src );
  SciGPU::HHV::HHVolumeGPU<long,2>::convertor dstConvert( dst, false );

  this->Inverter( dstConvert.get(), srcConvert.getConst() );
}


// -----------------------------------

__global__ void InvertKernel( long* dst,
                              const long* src,
                              const unsigned int nVals ) {
  const unsigned int idx = threadIdx.x + ( blockIdx.x * blockDim.x );
  
  const long invertVal = 255;

  if( idx < nVals ) {
    dst[idx] = invertVal - src[idx];
  }
}


void InvertImageGPU::Inverter( SciGPU::HHV::HHVolumeGPU<long,2>& dst,
                               const SciGPU::HHV::HHVolumeGPU<long,2>& src ) const {

  dst.Allocate( src.dimensions() );

  dim3 grid, threads;
  threads.x = 256;
  threads.y = threads.z = 1;
  grid.x = static_cast<unsigned int>( ceil( static_cast<float>(src.hyxels())/
                                            threads.x ) );

  
  InvertKernel<<<grid,threads>>>( dst.start(), src.start(), src.hyxels() );
  cudaError_t err = cudaThreadSynchronize();
  if( cudaSuccess != err ) {
    throw std::runtime_error( "InvertKernel failed" );
  }
}
